//------------------------头文件部分------------------------
#include <bits/stdc++.h>         // 包含所有C++标准头文件（实战中建议替换为具体需要的头文件）
#include <cuda.h>                // CUDA核心功能头文件（设备管理、上下文等）
#include "cuda_runtime.h"        // CUDA运行时API（内存管理、设备操作）
#include "device_launch_parameters.h" // 设备启动参数（blockIdx, threadIdx等）
#include <time.h>                // 基本时间函数（clock()等）
#include <sys/time.h>            // 高精度时间结构体（timeval）

//----------------------宏定义部分--------------------------
#define THREAD_PER_BLOCK 256     // 每个线程块包含256个线程（根据GPU架构调整优化）

//---------------------CUDA核函数部分-----------------------
__global__ void reduce0(float *d_in, float *d_out, unsigned int N) {
    // 声明共享内存（每个block创建独立副本）
    __shared__ float sdata[THREAD_PER_BLOCK];  // 存储块内中间结果

    // 线程索引计算
    unsigned int tid = threadIdx.x;            // 线程在块内的局部ID（0~255）

    // 计算当前block的全局内存起始地址
    float* inputbegin = d_in + blockIdx.x * blockDim.x;  // blockIdx.x为块索引

    // 将全局内存数据加载到共享内存（注意：缺少越界保护）
    sdata[tid] = inputbegin[tid];  // 当N非blockDim.x整数倍时可能越界
    __syncthreads(); //这里有没有 都正常
    // 归约计算阶段（低效的隔点累加模式）
    for(unsigned int s = 1; s < blockDim.x; s *= 2) {  // 步长指数增长
        // 条件判断导致线程束分化（部分线程不工作）
        if (tid % (2*s) == 0) {  // 仅选择特定间隔的线程参与计算
            sdata[tid] += sdata[tid + s];  // 跨距累加（s=1时相邻元素相加）
        }
        __syncthreads();         // 块内同步（确保本轮所有写入完成）
    }

    // 结果写回全局内存（仅首个线程执行）
    if (tid == 0) {               // 块内线程0负责输出
        d_out[blockIdx.x] = sdata[0];  // 将块归约结果写入对应位置
    }
}

//-------------------结果验证函数部分----------------------
bool check(float *out, float *res, int n) {
    for(int i=0; i<n; i++) {
        // 绝对误差检查（对大规模数据可能不够严谨）
        //if(abs(out[i]-res[i])>0.005) { // 建议改用相对误差：abs(a-b)/(abs(b)+eps)
        //    return false;        // 任意元素误差超限即返回错误
        float rel_err = abs(out[i]-res[i])/ (abs(res[i]) + 1e-5);
        if(rel_err > 0.1f) {
           return false;
        }
    }
    return true;                 // 所有元素在误差范围内
}

//---------------------主函数部分--------------------------
int main() {
    // 问题规模定义
    const int N = 32*1024*1024;  // 处理32M个浮点数（总数据量128MB）

    // 主机内存分配
    float *input = (float *)malloc(N*sizeof(float));    // 主机输入数组
    float *output = (float *)malloc((N/THREAD_PER_BLOCK)*sizeof(float)); // GPU输出数组
    float *result = (float *)malloc((N/THREAD_PER_BLOCK)*sizeof(float)); // CPU参考结果

    // 设备内存分配
    float *d_input, *d_output;
    cudaMalloc((void **)&d_input, N*sizeof(float));          // 设备端输入缓冲区
    cudaMalloc((void **)&d_output, (N/THREAD_PER_BLOCK)*sizeof(float)); // 设备端输出缓冲区

    // 检查设备内存分配结果
    if (!d_input || !d_output ) { 
        printf("GPU内存分配失败\n");  // 实际应使用cudaError_t获取具体错误信息
        return -1;
    }

    // 数据初始化（未设置随机种子导致每次运行结果相同）
    for(int i=0; i<N; i++) {
        input[i] = 2.0*(float)drand48()-1.0;  // 生成[-1,1)区间均匀分布随机数
    }

    // CPU参考计算（假设N是THREAD_PER_BLOCK的整数倍）
    int block_num = N/THREAD_PER_BLOCK;  // 总块数
    for(int i=0; i<block_num; i++) {
        float cur = 0;
        for(int j=0; j<THREAD_PER_BLOCK; j++) {  
            // 危险：当N非THREAD_PER_BLOCK整数倍时越界
            cur += input[i*THREAD_PER_BLOCK + j]; 
        }
        result[i] = cur;  // 存储每个块的累加结果
    }

    // 数据传输：主机->设备
    cudaError_t err = cudaMemcpy(d_input, input, N*sizeof(float), cudaMemcpyHostToDevice);
    if (err != cudaSuccess) {  // 检查数据传输错误
         printf("H2D Memcpy failed: %s\n", cudaGetErrorString(err));
    }

    // 核函数配置参数
    dim3 Grid(N/THREAD_PER_BLOCK, 1);  // 一维网格布局（每个block处理256元素）
    dim3 Block(THREAD_PER_BLOCK, 1);    // 一维线程块布局

    // 启动核函数（未检查内核启动错误）
    reduce0<<<Grid, Block>>>(d_input, d_output, N);  // 显式传递N参数

    // 数据传输：设备->主机
    cudaError_t err2 = cudaMemcpy(output, d_output, block_num*sizeof(float), cudaMemcpyDeviceToHost);
    if (err2 != cudaSuccess) {
         printf("D2H Memcpy failed: %s\n", cudaGetErrorString(err2));
    }

    // 验证GPU计算结果
    if(check(output, result, block_num)) {
        printf("验证通过：GPU与CPU结果一致\n");
    } else {
        printf("验证失败：结果不一致\n");
        // 打印错误结果（大数据量时建议限制打印数量）
        for(int i=0; i<block_num; i++) {
            printf("块%d: GPU=%lf vs CPU=%lf\n", i, output[i], result[i]);
        }
    }

    // 释放所有资源
    free(input);      // 释放主机输入内存
    free(output);     // 释放主机输出内存
    free(result);     // 释放CPU结果内存
    cudaFree(d_input);  // 释放设备输入内存
    cudaFree(d_output); // 释放设备输出内存
    
    return 0;  // 隐式返回（C++标准要求main函数返回int）
}